Fp16 constant weights - #3015
Conversation
Codecov Report❌ Patch coverage is
Additional details and impacted files@@ Coverage Diff @@
## main #3015 +/- ##
==========================================
- Coverage 83.72% 79.00% -4.73%
==========================================
Files 257 257
Lines 54822 54830 +8
Branches 4693 4695 +2
==========================================
- Hits 45901 43317 -2584
- Misses 8110 10627 +2517
- Partials 811 886 +75
Flags with carried forward coverage won't be shown. Click here to find out more. ☔ View full report in Codecov by Harness. 🚀 New features to boost your workflow:
|
The Weights of a non-expanded derivative were always built at the default precision, so a `float16` stencil got `float` coefficients. Every wavefield*weight product then bound to the mixed-precision operators and was promoted, defeating the point of the half-precision wavefield.
`_gen_value` printed the initializer with the printer's default dtype rather than the Array's, which stamped a `float` suffix onto the entries of a `double` Array and silently rounded them to single precision. Route it through a new `initvalue` printer hook, which also gives the targets a place to specialize an initializer whose type cannot be built from a plain literal.
`_prec` floors an untyped real literal at `float32` so that an integer default doesn't degrade the arithmetic around it. That floor also caught `float16`, which is never a fallback but an explicit request, so every literal in a half-precision Operator printed one type too wide. Only apply the floor when the default is not already a real type.
The stability check sums the whole field and asks whether the result is finite. The accumulator took the field's own dtype, so in half precision it overflowed within a few thousand points and reported an instability that wasn't there -- making `errctl=max`, the very option one reaches for to diagnose a suspected instability, unusable exactly where it is needed. Give it at least single precision.
A real literal in an otherwise integer expression is emitted at the Operator's precision, floored at `float32` so that an integer default does not degrade it. An Operator working in half wants that floor most of the time -- half is a storage format, and the accuracy of the literals is worth more than the width of the multiply -- but not always. Give the printer a flag for it, off by default, and have `_printer` pick up a Target's second printer where one is offered.
The same coefficients at two precisions are two different arrays, but neither `__eq__` nor `_hashable_content` looked at the dtype, so the first one built answered for both. An Operator asking for its weights in one precision would be handed whichever an earlier Operator had cached. Compare and hash on it. The name goes in rather than the type itself, which does not order and so cannot be sorted alongside the rest.
Whether an Operator working in half also computes in half decides what is calculated, not how quickly: the literals and the FD coefficients are rounded to three decimal digits. That is a mathematical choice, so it belongs with `interp-mode` in `sym_opt` rather than among the codegen options, and is validated and defaulted alongside it.
75d6b31 to
9759188
Compare
FabioLuporini
left a comment
There was a problem hiding this comment.
IIRC (a may be wrong and confusing it with the other PRO PR) there may have been leftover tiny comments in the old OSS PR that we might be able to address here
|
|
||
| INTERP_MODE = 'direct' | ||
|
|
||
| HALF_ARITH = False |
There was a problem hiding this comment.
you've got to move it down below -- currently it's right in between "INTERP_MODE" and its docstring
|
|
||
| def _hashable_content(self): | ||
| return (self.name, self.dimension, str(self.weights), self.scope) | ||
| # NOTE: `dtype` belongs here. The same coefficients at two precisions |
There was a problem hiding this comment.
this NOTE can go, it's obvious
| string, delegating to the printer so that languages whose types cannot | ||
| be built from plain literals (e.g. CUDA's `__half`) can specialize it. | ||
| """ | ||
| printer = get_printer(self.printer, obj.dtype) |
There was a problem hiding this comment.
this seems weird to me:
get_printer(self.printer
| be built from plain literals (e.g. CUDA's `__half`) can specialize it. | ||
| """ | ||
| printer = get_printer(self.printer, obj.dtype) | ||
| return printer.initvalue(init, obj.dtype) |
There was a problem hiding this comment.
why, instead of this new private method, don't you just pass self.ccode(obj) and let the printer handle the whole printing of the Array? ie u let it return c.Initializer(...)
| def _print_ListInitializer(self, expr): | ||
| return f"{{{', '.join(self._print(i) for i in expr.params)}}}" | ||
|
|
||
| def initvalue(self, init, dtype): |
There was a problem hiding this comment.
if this really is necessary (see other comments below) and if it really needs to be a public method, it should be put among the public methods, not among the private ones
fd41c9e to
9f3c81d
Compare
An Array initializer already reached the printer, via `ccode` on a `ListInitializer`; what it did not do was reach it at the Array's own precision. The elements were printed with the Operator's settings, so a narrow Array sitting in an Operator whose arithmetic is left at the default width had its entries emitted at the wider type. Pass the Array's dtype to `ccode`, as `Expression` already does, rather than route the initializer around the printer through an `initvalue` hook of its own. A target that needs to spell its literals differently overrides `_print_ListInitializer`, which is the ordinary extension point. With the precision now correct at the point of printing, `_prec` no longer needs to be told whether the arithmetic was narrowed on purpose: a real literal takes the precision it is being printed at, and the `float32` floor applies only where that is not itself a float. That is the same value as before for every dtype other than `float16`.
9f3c81d to
38dbac3
Compare
On top of #3012
TODO:
Add a knob for fp16 weights